﻿#include "cuda_runtime.h"
#include "cuda_profiler_api.h"

#include "helper_timer.h"
#include "Windows.h"

#include <stdlib.h>
#include <vector>
#include <algorithm>
#include "moderngpu.cuh"
#include "include/kernel.cuh"
#include "include/cudaEsg.h"
#include "include/esgSet.h"
#include "include/CudaComputeContext.h"
#include <memory>






namespace cudaEsg
{
	using namespace std;

	struct CompanyAttributesComparer : binary_function<esgKernels::CompanyAttributes, esgKernels::CompanyAttributes, bool>
	{
		__host__ __device__ __forceinline__ bool operator()(const esgKernels::CompanyAttributes& lhs, const esgKernels::CompanyAttributes& rhs) const
		{
			if (lhs.segment_index < rhs.segment_index)
				return true;
			if (rhs.segment_index < lhs.segment_index)
				return false;
			return lhs.score > rhs.score;
		}
	};


	/*template<typename T>
	struct greater : public std::binary_function<T, T, bool> {
		MGPU_HOST_DEVICE bool operator()(T a, T b) { return a > b; }
	};
*/
	operation_errors_t CudaComputeContext::RepeatGroupHeads(int groupsHeadsCount, int clusterCount, int companiesCount, int* groupsHeads, int& repeteadHeadsCount, int*& d_heads) const
	{
	
		int* h_heads;
		//int* d_heads = nullptr;
		repeteadHeadsCount = (groupsHeadsCount + 1) * clusterCount - 1;

		h_heads = static_cast<int*>(malloc(repeteadHeadsCount * sizeof(int)));

		if (h_heads == nullptr)
		{
			return operation_errors_t(esgMemoryAllocationError);
		}

		int currentIndex = 0;
		for (int i = 0; i < clusterCount; i++)
		{
			if (i > 0)
			{
				h_heads[currentIndex] = companiesCount * i;
				++currentIndex;
			}

			for (int j = 0; j < groupsHeadsCount; j++)
			{
				h_heads[currentIndex] = groupsHeads[j] + companiesCount * i;
				++currentIndex;
			}
		}

		cudaError_t cudaStatus = cudaMalloc((void **)&d_heads, repeteadHeadsCount * sizeof(int));

		if (cudaStatus != cudaSuccess)
		{
			return getOperationErrorFromCudaStatusMacro(cudaStatus);
		}

		cudaStatus = cudaMemcpy(d_heads, h_heads, repeteadHeadsCount * sizeof(int), cudaMemcpyHostToDevice);
		free(h_heads);
		if (cudaStatus != cudaSuccess)
		{
			return getOperationErrorFromCudaStatusMacro(cudaStatus);
		}
		return operation_errors_t::success();

		//set->groupHeadsCount = repeteadHeadsCount;
	}

	operation_errors_t CudaComputeContext::Ratings(int setId, int universeId, int weightMatrixId, unsigned char filter, int groupHeadsCount, int* groupsHeads, RatingsResult& result)
	{
		cudaProfilerStart();
		cudaError_t cudaStatus;

		cudaGetLastError(); //reset errors;

		this->ResetStatistics();
		ESGSet* set = GetSet(setId);

		if (set == nullptr)
		{
			return operation_errors_t(esgSetNotFound);
		}

		if (set->companiesCount != result.companiesCount)
			return operation_errors_t(esgCompaniesCountMismatch, "");

		if (set->indicatorsCount != result.indicatorsCount)
			return operation_errors_t(esgIndicatorsCountMismatch);


		float* v_d = set->indicators_scores;
		if (v_d == nullptr)
			return operation_errors_t(esgIndicatorsScoresNotFound);

		float* p_d = set->GetMatrix(weightMatrixId);
		if (p_d == nullptr)
			return operation_errors_t(esgMatrixNotFound);


		unsigned char* flags_d = set->GetProfiles(weightMatrixId);
		if (flags_d == nullptr)
			return operation_errors_t(esgMatrixInfoNotFound);


		unsigned char* u_d = set->GetUniverse(universeId);
		if (u_d == nullptr)
			return operation_errors_t(esgUniverseNotFound);

		unsigned char* attributes_d = set->attributes;
		if (attributes_d == nullptr)
			return operation_errors_t(esgAttributesNotFound);

		sdkResetTimer(&_timer);
		sdkStartTimer(&_timer);

		int companiesNo = set->companiesCount;
		int indicatorsNo = set->indicatorsCount;
		byte clustersCount = set->CLUSTERS_COUNT;
		byte disjunctClustersCount = set->DISJUNCT_CLUSTERS_COUNT;
		int ratingsCount = clustersCount * companiesNo;



		try
		{
			_mgpuContext->Start();


			cudaStatus = cudaMemcpyToSymbol(esgEnds, set->h_configEgsHeads, sizeof(byte) * set->h_configEsgHeadsCount);
			returnOnErrorMacro(cudaStatus);
			cudaStatus = cudaMemcpyToSymbol(clusterStarts, set->h_configDisjunctClustersHeads, sizeof(byte) * set->h_configDisjunctClustersHeadsCount);
			returnOnErrorMacro(cudaStatus);
			cudaStatus = cudaMemcpyToSymbol(clusterGroups, set->h_configDisjunctClustersMap, sizeof(byte) * set->h_configDisjunctClustersMapCount);
			returnOnErrorMacro(cudaStatus);


			int* d_groupHeads;
			int repeatedGroupHeads;
			RepeatGroupHeads(groupHeadsCount, clustersCount, companiesNo, groupsHeads, repeatedGroupHeads, d_groupHeads);


			float* h_ratingsPinned;
			cudaStatus = cudaMallocHost((void**)&h_ratingsPinned, sizeof(float) * ratingsCount);
		
			returnOnErrorMacro(cudaStatus);
			result.timedOperations->push_back(make_unique<Operation>("cudaMallocHost<h_ratings>", _mgpuContext->Split()));
			this->AddMemoryOperation("host - h_ratingsPinned", sizeof(float) * ratingsCount);

			if ((result.options & DefaultWeightMatrix) != 0)
			{
				cudaStatus = cudaMemcpy(h_ratingsPinned, set->scores, ratingsCount * sizeof(float), cudaMemcpyDeviceToHost);
				returnOnErrorMacro(cudaStatus);
			}



			int* h_companies_map;
			cudaStatus = cudaMallocHost((void**)&h_companies_map, sizeof(int) * ratingsCount);//throws CudaException
			returnOnErrorMacro(cudaStatus);
			result.timedOperations->push_back(make_unique<Operation>("cuda_mallocHost<h_companies_map>", _mgpuContext->Split()));
			this->AddMemoryOperation("host - h_companies_map", sizeof(int) * ratingsCount);

			if ((result.options & DefaultWeightMatrix) != 0)
			{
				int* companiesIndices = static_cast<int*>(malloc(sizeof(int)*ratingsCount));
				for (int line = 0; line < companiesNo; line++)
				{
					for (int i = 0; i < clustersCount; i++)
					{
						companiesIndices[i * companiesNo + line] = line;
					}
				}
				memcpy(h_companies_map, companiesIndices, ratingsCount * sizeof(int));
			}



			float* h_indicatorsRatings;
			float* h_reference_indicatorsRatings = nullptr;
			if ((result.options & ComputeIndicatorsRatingsOption) != 0)
			{
				cudaMallocHost((void**)&h_indicatorsRatings, sizeof(float) * companiesNo * indicatorsNo);//throws CudaException
				result.timedOperations->push_back(make_unique<Operation>("cuda_mallocHost<h_indicatorsRatings>", _mgpuContext->Split()));

				h_reference_indicatorsRatings = h_indicatorsRatings;
				this->AddMemoryOperation("host - h_indicatorsRatings", companiesNo * indicatorsNo * sizeof(float));
			}

			float* h_aggregatedWeights;
			float* h_reference_aggregatedWeights = nullptr;
			if (((result.options & ComputeAggregatedWeights) != 0)&&((result.options & DefaultWeightMatrix)==0))
			{
				cudaMallocHost((void**)&h_aggregatedWeights, sizeof(float) * ratingsCount);//throws CudaException
				result.timedOperations->push_back(make_unique<Operation>("cuda_mallocHost<h_aggregatedWeights>", _mgpuContext->Split()));
				h_reference_aggregatedWeights = h_aggregatedWeights;
				this->AddMemoryOperation("host - aggragatedWeights", ratingsCount * sizeof(float));
			}

			if ((result.options & DefaultWeightMatrix) == 0)
			{
				esgKernels::ComputeRatings((set->h_configEsgHeadsCount + 1), v_d, p_d, flags_d, attributes_d, filter, h_ratingsPinned, h_companies_map, h_reference_indicatorsRatings, h_reference_aggregatedWeights, companiesNo, indicatorsNo, disjunctClustersCount, clustersCount);

				cudaStatus = cudaPeekAtLastError(); //launch errors;
				returnOnErrorMacro(cudaStatus);
				result.timedOperations->push_back(make_unique<Operation>("kernel<computeRatingsKernel>", _mgpuContext->Split()));
			}

			mgpu::SegSortPairsFromIndices(h_ratingsPinned, h_companies_map, ratingsCount, d_groupHeads, repeatedGroupHeads, *_mgpuContext, mgpu::greater<float>());

			result.timedOperations->push_back(make_unique<Operation>("kernelSegment<sort>", _mgpuContext->Split()));

			memcpy(result.ratings, h_ratingsPinned, ratingsCount * sizeof(float));
			result.timedOperations->push_back(make_unique<Operation>("toHostMemory<d_ratings>", _mgpuContext->Split()));

			memcpy(result.indices, h_companies_map, ratingsCount * sizeof(int));
			result.timedOperations->push_back(make_unique<Operation>("toHostMemory<d_companies_map>", _mgpuContext->Split()));

			if ((result.options & ComputeIndicatorsRatingsOption) != 0)
			{
				memcpy(result.indicators_ratings, h_indicatorsRatings, companiesNo * indicatorsNo * sizeof(float));

				cudaFreeHost(h_indicatorsRatings);
				returnOnErrorMacro(cudaStatus);
				result.timedOperations->push_back(make_unique<Operation>("toHostMemory<h_indicatorsRatings>", _mgpuContext->Split()));
			}


			if ((result.options & ComputeAggregatedWeights) != 0)
			{
				if ((result.options & DefaultWeightMatrix) != 0)
				{
					cudaStatus = cudaMemcpy(result.aggregatedWeights, set->aggregatedWeights, ratingsCount * sizeof(float), cudaMemcpyDeviceToHost);
					returnOnErrorMacro(cudaStatus);
				}
				else
				{
					memcpy(result.aggregatedWeights, h_aggregatedWeights, ratingsCount * sizeof(float));

					result.timedOperations->push_back(make_unique<Operation>("toHostMemory<h_aggragatedWeights>", _mgpuContext->Split()));
					cudaFreeHost(h_aggregatedWeights);
				}
			}


			if ((result.options & ReturnWeights) != 0 && (result.weights != nullptr))
			{
				cudaStatus = cudaMemcpy(result.weights, p_d, companiesNo * indicatorsNo * sizeof(float), cudaMemcpyDeviceToHost);
				if (cudaStatus != cudaSuccess)
				{
					returnOnErrorMacro(cudaStatus);
				}
				result.timedOperations->push_back(make_unique<Operation>("toHostMemory<Weights>", _mgpuContext->Split()));
			}

			sdkStopTimer(&_timer);
			float ratingsMilliseconds = sdkGetTimerValue(&_timer);

			sdkResetTimer(&_timer);
			sdkStartTimer(&_timer);

		

			ComputeRanking(result, h_ratingsPinned, set, groupHeadsCount, groupsHeads, u_d, h_companies_map,filter);
			sdkStopTimer(&_timer);

			float rankingsMilliseconds = sdkGetTimerValue(&_timer);
			result.timedOperations->push_back(make_unique<Operation>("ranking", rankingsMilliseconds / 1000));

			float totalTimeMilliseconds = ratingsMilliseconds + rankingsMilliseconds;
			result.timedOperations->push_back(make_unique<Operation>("TotalRankingTime", totalTimeMilliseconds / 1000));

			cudaFree(d_groupHeads);
			cudaFreeHost(h_ratingsPinned);
			cudaFreeHost(h_companies_map);
			cudaMemset(clusterStarts, 0, sizeof(int) * MAX_DISJUNCT_CLUSTERS_HEADS);
			cudaMemset(clusterGroups, 0, sizeof(int) * MAX_DISJUNCT_CLUSTERS_GROUPS_MAP_COUNT);
		}
		// ReSharper disable once CppEntityNeverUsed
		catch (mgpu::MgpuArchVersionException& exception)
		{
			return operation_errors_t(esgCudaInvalidArch, "MgpuArchVersion: this CUDA executable was not compiled with support for current device arch");
		}
		catch (mgpu::CudaException& exception)
		{
			cudaError_t status = exception.error;
			returnOnErrorMacro(status);
		}
		catch (...)
		{
			return operation_errors_t(esgError);
		}

		cudaProfilerStop();
		return operation_errors_t::success();
	}

	void CudaComputeContext::ComputeRanking(RatingsResult& result, float* ratings_h, ESGSet* set, int clusterGroupHeadsCount, int* clusterGroupHeads, unsigned char* universe, int* companiesMap, unsigned char filter)
	{
		int* rankings = result.rankings;
		float* mins = result.group_mins;
		float* maxs = result.group_maxs;
		float* avgs = result.group_avgs;
		int* groupsCounts = result.group_counts;

		float* percentiles = result.percentile;
		bool* isInReferenceUniverse = result.isInReferenceUniverse;

		int companiesNo = result.companiesCount;

		int* segmentedGroupCount = (int*)calloc((clusterGroupHeadsCount + 1), sizeof(int));


		int nextChangeIndex = 0;
		int currentHeadIndex = -1;

		int rank = 0;

		for (int i = 0; i <= companiesNo; i++)
		{
			if (i == nextChangeIndex) //head group change
			{
				if (i != 0)
				{
					segmentedGroupCount[currentHeadIndex] = rank;
				}

				if (i == companiesNo)
				{
					break;
				}
				rank = 0;
				nextChangeIndex = companiesNo;
				++currentHeadIndex;
				if (currentHeadIndex < clusterGroupHeadsCount)
				{
					nextChangeIndex = clusterGroupHeads[currentHeadIndex];
				}
			}

			if (universe[companiesMap[i]] == 1 && ratings_h[i] >= -3 && ((set->hostAttributes[companiesMap[i]]& filter) == filter || (!filter)))
			{
				++rank;
				isInReferenceUniverse[companiesMap[i]] = true;
			}
			else
			{
				isInReferenceUniverse[companiesMap[i]] = false;
			}
		}

		float* currentMin = mins;
		float* currentMax = maxs;
		float* currentAvg = avgs;
		int* currentGroupCount = groupsCounts;
		int clustersCount = set->CLUSTERS_COUNT;

		for (int clusterIndex = 0; clusterIndex < clustersCount; clusterIndex++)
		{
			int baseIndex = clusterIndex * companiesNo;
			nextChangeIndex = 0;
			currentHeadIndex = -1;
			float sum = 0;
			rank = 0;
			int computedRank = 1;

			*currentMin = -1;
			*currentMax = -1;
			*currentAvg = -1;

			for (int i = 0; i <= companiesNo; i++)
			{
				if (i == nextChangeIndex) //head group change
				{
					if (i != 0)
					{
						if (rank > 0)
						{
							*currentAvg = sum / rank;
						}
						else
						{
							*currentAvg = -1;
						}

						++currentAvg;
						++currentMax;
						++currentMin;

						*currentMin = -1;
						*currentMax = -1;
						*currentAvg = -1;
					}

					if (i == companiesNo)
						break;

					sum = 0;
					rank = 0;
					computedRank = 1;

					nextChangeIndex = companiesNo;
					++currentHeadIndex;

					if (currentHeadIndex < clusterGroupHeadsCount)
					{
						nextChangeIndex = clusterGroupHeads[currentHeadIndex];
					}
				}

				if (i == companiesNo)
					break;

				int currentIndex = baseIndex + i;
				float rating = ratings_h[currentIndex];

				if (rating <= -3 || !((set->hostAttributes[companiesMap[currentIndex]] & filter) == filter || (!filter)))
				{
					rankings[currentIndex] = -1;

					ratings_h[currentIndex] = (-1) * (rating + 4); //reverse filter transformation

					continue;
				}
				else if (rating != rating)
				{ //NAN number
					rankings[currentIndex] = -2;

					continue;
				}

				if (universe[companiesMap[currentIndex]])
				{
					++rank;

					if (*currentMax == -1)
					{
						*currentMax = rating;
					}

					*currentMin = rating;
					sum += rating;
				}

				if (i == 0 || (rating != ratings_h[currentIndex - 1]))
				{
					if (universe[companiesMap[currentIndex]])
					{
						computedRank = rank;
					}
					else
					{
						computedRank = rank + 1;
					}
				}

				if (universe[companiesMap[currentIndex]])
				{
					currentGroupCount[currentIndex] = segmentedGroupCount[currentHeadIndex];
				}
				else
				{
					currentGroupCount[currentIndex] = segmentedGroupCount[currentHeadIndex] + 1;
				}
				rankings[currentIndex] = computedRank;
				percentiles[currentIndex] = esgKernels::ComputePercentile(currentGroupCount[currentIndex], computedRank);
			}

			//percentile = ((float)(groupCount - rank)) / (groupCount - 1);
		}
		free(segmentedGroupCount);
	}

	CudaComputeContext::CudaComputeContext()
	{
		try
		{
			_mgpuContext = mgpu::CreateCudaDevice(0, nullptr, true);
		}
		// ReSharper disable once CppEntityNeverUsed
		catch (mgpu::MgpuArchVersionException& exception)
		{
			operation_errors_t operation = operation_errors_t(esgCudaInvalidArch, "MgpuArchVersion: this CUDA executable was not compiled with support for current device arch");
			throw ComputeContextException(operation);
		}
		catch (mgpu::CudaException& exception)
		{
			cudaError_t status = exception.error;
			operation_errors_t operation = operation_errors_t(esgCudaError, status, cudaGetErrorString(status), "", -1, "");
			throw ComputeContextException(operation);
		}
		catch (...)
		{
			throw ComputeContextException(operation_errors_t(esgError));
		}
		sets = new vector<tuple<int, ESGSet*>>();
		memoryStatistics = new MemoryOperations();
		sdkCreateTimer(&_timer);
	}

	CudaComputeContext::~CudaComputeContext()
	{
		delete _timer;
		delete memoryStatistics;

		int* ids;
		int n = 0;
		CudaComputeContext::CleanContext(ids, n);
		delete sets;
		sets = nullptr;
	}

	operation_errors_t CudaComputeContext::CleanContext(int* & ids, int& count)
	{
		cudaGetLastError(); //reset errors;
		operation_errors_t returnStatus = operation_errors_t::success();

		if (count == 0)
		{
			count = sets->size();
			ids = static_cast<int*>(calloc(count + 1, sizeof(int)));
		}
		int i = 0;
		while (!sets->empty())
		{
			ESGSet* set = get<1>(sets->back());
			if (i < count)
			{
				ids[i] = set->id;
			}
			++i;
			cudaEsg_errors_t cudaStatus = CleanSet(set);
			if (cudaStatus != cudaSuccess)
			{
				returnStatus = operation_errors_t(cudaStatus);
				break;
			}

			sets->pop_back();
		}
		sets->clear();
		return returnStatus;
	}

	operation_errors_t CudaComputeContext::AddSet(SetConfiguration* configuration)
	{

		cudaGetLastError(); //reset errors;

		if (configuration->companiesCount <= 0)
			return operation_errors_t(esgInvalidParameters);

		if (configuration->esgHeadsCount > ESG_ENDS_COUNT)
		{
			return operation_errors_t(esgInvalidEsgEndsCount, "There must be at least one theme");
		}


		if (configuration->setId>=0)
		{
			return operation_errors_t(esgInvalidSetId, "Set Id is internally generated");
		}

		/*	ESGSet* set = GetSet(configuration->setId);

		if (set != nullptr)
			return operation_errors_t(esgDuplicateEntity);*/

		if (configuration->DISJUNCT_CLUSTERS_COUNT > MAX_DISJUNCT_CLUSTERS_COUNT)
		{
			return operation_errors_t(esgMaxDisjunctClustersCount);
		}

		if (configuration->disjunctClustersHeadsCount > MAX_DISJUNCT_CLUSTERS_HEADS)
		{
			return operation_errors_t(esgMaxDisjunctClustersHeads);
		}

		configuration->setId = this->NextId();

		//	int CLUSTERS_COUNT = configuration->CLUSTERS_COUNT;

		ESGSet* set = new ESGSet(configuration->setId, configuration->companiesCount);

		set->companiesCount = configuration->companiesCount;

		//	int* d_heads = nullptr;

		operation_errors_t returnStatus = operation_errors_t::success();

		{
			set->h_configEgsHeads = static_cast<byte*>(malloc(sizeof(byte) * configuration->esgHeadsCount));
			memcpy(set->h_configEgsHeads, configuration->esgHeads, configuration->esgHeadsCount * sizeof(byte));
			set->h_configEsgHeadsCount = configuration->esgHeadsCount;

			this->AddMemoryOperation("host - h_configEgsHeads", _msize(set->h_configEgsHeads));

			set->h_configDisjunctClustersHeads = static_cast<byte*>(malloc(sizeof(byte) * configuration->disjunctClustersHeadsCount));
			memcpy(set->h_configDisjunctClustersHeads, configuration->disjunctClustersHeads, configuration->disjunctClustersHeadsCount * sizeof(byte));
			set->h_configDisjunctClustersHeadsCount = configuration->disjunctClustersHeadsCount;

			this->AddMemoryOperation("host - h_configDisjunctClustersHeads", _msize(set->h_configEgsHeads));

			set->h_configDisjunctClustersMap = static_cast<byte*>(malloc(sizeof(byte) * configuration->disjunctClustersMapCount));
			memcpy(set->h_configDisjunctClustersMap, configuration->disjunctClustersMap, configuration->disjunctClustersMapCount * sizeof(byte));
			set->h_configDisjunctClustersMapCount = configuration->disjunctClustersMapCount;
			this->AddMemoryOperation("host - h_configDisjunctClustersMap", _msize(set->h_configDisjunctClustersMap));

			set->indicatorsCount = configuration->indicatorsCount;
			set->CLUSTERS_COUNT = configuration->CLUSTERS_COUNT;
			set->DISJUNCT_CLUSTERS_COUNT = configuration->DISJUNCT_CLUSTERS_COUNT;
		}
		if (returnStatus.esgError != esgSuccess)
		{
			CleanSet(set);
			delete set;
			return returnStatus;
		}

		sets->push_back(make_tuple(configuration->setId, set));

		return esgSuccess;
	}

	operation_errors_t CudaComputeContext::SetIndicators(int setId, int companiesCount, int indicatorsCount, float* indicatorScores)
	{
		cudaGetLastError(); //reset errors;

		if (indicatorScores == nullptr || indicatorScores <= 0 || companiesCount <= 0)
			return operation_errors_t(esgInvalidParameters);

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (set->companiesCount != companiesCount)
			return operation_errors_t(esgParametersMismatch);

		if (set->indicatorsCount != indicatorsCount)
			return operation_errors_t(esgParametersMismatch);

		int matrixElementCount = companiesCount * indicatorsCount;

		float* d_indicators = nullptr;

		operation_errors_t returnStatus = operation_errors_t::success();

		for (int i = 0; i < 1; i++)
		{
			cudaError_t cudaStatus = cudaMalloc((void **)&d_indicators, matrixElementCount * sizeof(float));

			this->AddMemoryOperation("cuda - indicators", matrixElementCount * sizeof(float));

			if (cudaStatus != cudaSuccess)
			{
				returnStatus = getOperationErrorFromCudaStatusMacro(cudaStatus);
				break;
			}

			cudaStatus = cudaMemcpy(d_indicators, indicatorScores, matrixElementCount * sizeof(float), cudaMemcpyHostToDevice);
			if (cudaStatus != cudaSuccess)
			{
				returnStatus = getOperationErrorFromCudaStatusMacro(cudaStatus);
				break;
			}
		}

		if (returnStatus.esgError != esgSuccess)
		{
			if (d_indicators != nullptr)
			{
				cudaFree(d_indicators);
			}

			return returnStatus;
		}

		float* previous = set->indicators_scores;
		if (previous != nullptr)
		{
			cudaFree(previous);
			//returnOnErrorMacro(cudaStatus);//do nothing if error
		}

		set->indicators_scores = d_indicators;

		return esgSuccess;
	}

	cudaEsg_errors_t CudaComputeContext::CleanSet(ESGSet* set) const
	{
		if (set == nullptr)
			return esgSuccess;

		try
		{
			if (set->indicators_scores != nullptr)
			{
				cudaFree(set->indicators_scores);
				set->indicators_scores = nullptr;
			}

			//if (set->d_groupHeads != nullptr){
			//	cudaFree(set->d_groupHeads);
			//	set->d_groupHeads = nullptr;
			//}

			if (set->attributes != nullptr)
			{
				cudaFree(set->attributes);
				set->attributes = nullptr;
			}

			if (set->universes != nullptr)
			{
				vector<tuple<int, unsigned char*>>* universes = set->universes;
				while (!universes->empty())
				{
					free(get<1>(universes->back()));
					universes->pop_back();
				}

				set->universes = nullptr;
				delete universes;
			}

			if (set->weightsMatrices != nullptr)
			{
				vector<tuple<int, float*>>* weightsMatrices = set->weightsMatrices;
				while (!weightsMatrices->empty())
				{
					cudaFree(get<1>(weightsMatrices->back()));
					weightsMatrices->pop_back();
				}

				set->weightsMatrices = nullptr;
				delete weightsMatrices;
			}


			if (set->profiles != nullptr)
			{
				vector<tuple<int, unsigned char*>>* profiles = set->profiles;
				while (!profiles->empty())
				{
					cudaFree(get<1>(profiles->back()));
					profiles->pop_back();
				}

				set->profiles = nullptr;
				delete profiles;
			}

			if (set->modifiers != nullptr)
			{
				vector<tuple<int, float*>>* modifiers = set->modifiers;
				while (!modifiers->empty())
				{
					cudaFree(get<1>(modifiers->back()));
					modifiers->pop_back();
				}

				set->modifiers = nullptr;
				delete modifiers;
			}

			if (set->profiles != nullptr)
			{
				vector<tuple<int, float*>>* issueWeights = set->issueWeights;
				while (!issueWeights->empty())
				{
					cudaFree(get<1>(issueWeights->back()));
					issueWeights->pop_back();
				}

				set->issueWeights = nullptr;
				delete issueWeights;
			}

			if (set->company_flags!=nullptr)
			{
				cudaFree(set->company_flags);
				set->company_flags = nullptr;
			}
			
			if (set->company_flag_types!=nullptr)
			{
				cudaFree(set->company_flag_types);
				set->company_flag_types = nullptr;
			}

			return esgSuccess;
		}
		catch (...)
		{
			return cudaEsg_errors_t(esgError);
		}
	}

	operation_errors_t CudaComputeContext::AddUniverse(int setId, int universeId, int companiesCount, unsigned char* universe)
	{
		cudaGetLastError(); //reset errors;
		if (companiesCount <= 0 || universe == nullptr)
			return operation_errors_t(esgInvalidParameters);

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (companiesCount != (set->companiesCount))
			return operation_errors_t(esgParametersMismatch);

		if (set->GetUniverse(universeId) != nullptr)
			return operation_errors_t(esgDuplicateEntity);


		unsigned char* h_universe;
		h_universe = static_cast<unsigned char*>(malloc(companiesCount * sizeof(unsigned char)));
		memcpy(h_universe, universe, companiesCount * sizeof(unsigned char));


		set->AddUniverse(universeId, h_universe);

		return operation_errors_t::success();
	}

	operation_errors_t CudaComputeContext::DeleteUniverse(int setId, int universeId)
	{
		cudaGetLastError(); //reset errors;

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (set->GetUniverse(universeId) == nullptr)
			return operation_errors_t(esgUniverseNotFound);

		return operation_errors_t::success();
	}

	operation_errors_t CudaComputeContext::SetAttributes(int setId, int companiesCount, unsigned char* attributes)
	{
		cudaGetLastError(); //reset errors;

		if (companiesCount <= 0 || attributes == nullptr)
			return esgInvalidParameters;

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (set->companiesCount % companiesCount != 0)
			return operation_errors_t(esgInvalidParameters);

		unsigned char* d_attributes = nullptr;

		cudaError_t cudaStatus = cudaMalloc((void **)&d_attributes, companiesCount * sizeof(unsigned char));
		returnOnErrorMacro(cudaStatus);


		set->hostAttributes = static_cast<unsigned char*>(malloc(sizeof(unsigned char) *companiesCount));
		memcpy (set->hostAttributes, attributes, companiesCount * sizeof(unsigned char));
	

		cudaStatus = cudaMemcpy(d_attributes, attributes, companiesCount * sizeof(unsigned char), cudaMemcpyHostToDevice);
		if (cudaStatus != cudaSuccess)
		{
			cudaFree(attributes);
			returnOnErrorMacro(cudaStatus);
		}

		unsigned char* previousAttributes = set->attributes;
		set->attributes = d_attributes;

		if (previousAttributes != nullptr)
		{
			cudaFree(previousAttributes);
			//returnOnErrorMacro(cudaStatus);//do nothing if error
		}

		return esgSuccess;
	}

	operation_errors_t CudaComputeContext::SetCompaniesScores(int setId, int scoresCount, float* scores)
	{
		cudaGetLastError(); //reset errors;

		if (scoresCount <= 0 || scores == nullptr)
			return esgInvalidParameters;

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (scoresCount % set->companiesCount != 0)
			return operation_errors_t(esgInvalidParameters);

		float* d_scores = nullptr;

		cudaError_t cudaStatus = cudaMalloc((void **)&d_scores, scoresCount * sizeof(float));
		returnOnErrorMacro(cudaStatus);

		cudaStatus = cudaMemcpy(d_scores, scores, scoresCount * sizeof(float), cudaMemcpyHostToDevice);
		if (cudaStatus != cudaSuccess)
		{
			cudaFree(scores);
			returnOnErrorMacro(cudaStatus);
		}

		float* previousAttributes = set->scores;
		set->scores = d_scores;

		if (previousAttributes != nullptr)
		{
			cudaFree(previousAttributes);
			//returnOnErrorMacro(cudaStatus);//do nothing if error
		}

		return esgSuccess;
	}

	operation_errors_t CudaComputeContext::SetCompaniesWeights(int setId, int scoresCount, float* scores)
	{
		cudaGetLastError(); //reset errors;

		if (scoresCount <= 0 || scores == nullptr)
			return esgInvalidParameters;

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (scoresCount % set->companiesCount != 0)
			return operation_errors_t(esgInvalidParameters);

		float* d_scores = nullptr;

		cudaError_t cudaStatus = cudaMalloc((void **)&d_scores, scoresCount * sizeof(float));
		returnOnErrorMacro(cudaStatus);

		cudaStatus = cudaMemcpy(d_scores, scores, scoresCount * sizeof(float), cudaMemcpyHostToDevice);
		if (cudaStatus != cudaSuccess)
		{
			cudaFree(scores);
			returnOnErrorMacro(cudaStatus);
		}

		float* previousAttributes = set->aggregatedWeights;
		set->aggregatedWeights = d_scores;

		if (previousAttributes != nullptr)
		{
			cudaFree(previousAttributes);
			//returnOnErrorMacro(cudaStatus);//do nothing if error
		}

		return esgSuccess;
	}

	//operation_errors_t CudaComputeContext::SetCompanyGroupFlags(int setId, int companiesCount, int attributes_Types_Count, vector<ComparativeGroup> attributes) 
	operation_errors_t CudaComputeContext::SetCompanyGroupFlags(int setId, int companiesCount, int attributes_Types_Count,int regionGroupsCount,int peerGroupsCount, int* attributes, int attributesCount)
	{
		cudaGetLastError(); //reset errors;

		if (companiesCount <= 0 )
			return esgInvalidParameters;

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (set->companiesCount != companiesCount )
			return operation_errors_t(esgInvalidParameters);

		set->flags_types_count = attributes_Types_Count;


		int* d_attributes = nullptr;

		cudaError_t cudaStatus = cudaMalloc((void **)&d_attributes, attributesCount * sizeof(int));
		returnOnErrorMacro(cudaStatus);
		cudaStatus = cudaMemcpy(d_attributes, attributes, attributesCount * sizeof(int), cudaMemcpyHostToDevice);
		
		int* company_flag_types = static_cast<int*>(malloc(sizeof(int)*attributes_Types_Count));

		company_flag_types[0] = regionGroupsCount; //TODO Madalina:refactor this;
		company_flag_types[1] = peerGroupsCount;

		int* d_attributes_type_count = nullptr;
		cudaStatus = cudaMalloc((void **)&d_attributes_type_count, attributes_Types_Count * sizeof(int));
		returnOnErrorMacro(cudaStatus);
		cudaStatus = cudaMemcpy(d_attributes_type_count, company_flag_types, attributes_Types_Count * sizeof(int), cudaMemcpyHostToDevice);


		/*int index_i = 0;
		for (auto collection : attributes)
		{
			set->company_flag_types[index_i] = collection.flag_type_count;
			int offset = index_i*companiesCount*sizeof(int);
			memcpy(mergedAttributes + offset, collection.flags, companiesCount*sizeof(int));
		}

		cudaStatus = cudaMemcpy(d_attributes, mergedAttributes, attributesCount * sizeof(int), cudaMemcpyHostToDevice);
*/
		if (cudaStatus != cudaSuccess)
		{
			cudaFree(d_attributes);
			returnOnErrorMacro(cudaStatus);
		}

		int* previousAttributes = set->company_flags;
		set->company_flags = d_attributes;
		set->company_flag_types = d_attributes_type_count;
		if (previousAttributes != nullptr)
		{
			cudaFree(previousAttributes);
			//returnOnErrorMacro(cudaStatus);//do nothing if error
		}
		
		return esgSuccess;
	}

	operation_errors_t CudaComputeContext::AddMatrix(int setId, int matrixId, int elementsCount, float* weightsMatrix, int flagsCount, unsigned char* profileFlags)
	{
		cudaGetLastError(); //reset errors;
		if (elementsCount <= 0 || weightsMatrix == nullptr)
			return operation_errors_t(esgInvalidParameters);


		if (flagsCount <= 0 || profileFlags == nullptr)
			return operation_errors_t(esgInvalidParameters);

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (elementsCount != (set->companiesCount) * (set->indicatorsCount))
			return operation_errors_t(esgParametersMismatch);

		if (flagsCount != (set->companiesCount))
			return operation_errors_t(esgParametersMismatch);

		if (set->GetMatrix(matrixId) != nullptr)
			return operation_errors_t(esgDuplicateEntity);

		float* d_matrix = nullptr;
		unsigned char* d_profiles = nullptr;

		cudaError_t cudaStatusWeights = cudaMalloc((void **)&d_matrix, elementsCount * sizeof(float));

		this->AddMemoryOperation("cuda - matrix", elementsCount * sizeof(float));
		returnOnErrorMacro(cudaStatusWeights);

		cudaStatusWeights = cudaMemcpy(d_matrix, weightsMatrix, elementsCount * sizeof(float), cudaMemcpyHostToDevice);
		if (cudaStatusWeights != cudaSuccess)
		{
			cudaFree(d_matrix);
			returnOnErrorMacro(cudaStatusWeights);
		}

		cudaError_t cudaStatusFlags = cudaMalloc((void **)&d_profiles, flagsCount * sizeof(unsigned char));

		if (cudaStatusFlags != cudaSuccess)
		{
			cudaFree(d_matrix);
			returnOnErrorMacro(cudaStatusFlags);
		}


		this->AddMemoryOperation("cuda - profiles", flagsCount * sizeof(unsigned char));

		cudaStatusFlags = cudaMemcpy(d_profiles, profileFlags, flagsCount * sizeof(unsigned char), cudaMemcpyHostToDevice);
		if (cudaStatusFlags != cudaSuccess)
		{
			cudaFree(d_matrix);
			cudaFree(d_profiles);
			returnOnErrorMacro(cudaStatusFlags);
		}

		set->AddMatrix(matrixId, d_matrix, d_profiles);

		return operation_errors_t::success();
	}

	operation_errors_t CudaComputeContext::DeleteMatrix(int setId, int matrixId)
	{
		cudaGetLastError(); //reset errors;

		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (set->GetMatrix(matrixId) == nullptr)
			return operation_errors_t(esgMatrixNotFound);

		return operation_errors_t::success();
	}

	operation_errors_t CudaComputeContext::DeleteSet(int setId)
	{
		try
		{
			ESGSet* setToDelete = GetSet(setId);
			CleanSet(setToDelete);

			sets->erase(std::remove_if(sets->begin(), sets->end(),
				[setId](tuple<int, ESGSet*> const& o)
			{
				return get<0>(o) == setId;
			}), sets->end());

			//sets.erase(std::remove_if(sets.begin(), sets.end(),
			//[setId]( const ESGSet &o){
			//return o.id == setId;
			//}), sets.end());

			return operation_errors_t::success();
		}
		catch (...)
		{
			return operation_errors_t(esgCudaError);
		}
	}

	operation_errors_t CudaComputeContext::FillSetAttributes(int setId, int& companiesCount, int& indicatorsCount, int& clusterGroupsCount, int& clustersCount)
	{
		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		companiesCount = set->companiesCount;
		indicatorsCount = set->indicatorsCount;
		clusterGroupsCount = set->clusterGroupHeadsCount + 1;
		clustersCount = set->CLUSTERS_COUNT;


		return operation_errors_t::success();
	}

	operation_errors_t CudaComputeContext::FillSetAttributes(int setId, int& companiesCount, int& indicatorsCount, int& clusterGroupsCount, int& clustersCount, int& company_grouping_types_count)
	{
		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		companiesCount = set->companiesCount;
		indicatorsCount = set->indicatorsCount;
		clusterGroupsCount = set->clusterGroupHeadsCount + 1;
		clustersCount = set->CLUSTERS_COUNT;
		company_grouping_types_count = set->flags_types_count;

		return operation_errors_t::success();
	}

	void CudaComputeContext::AddMemoryOperation(string name, double elapsedSeconds) const
	{
		memoryStatistics->operationsList->push_back(make_unique<MemoryOperation>(name, elapsedSeconds));
	}

	int CudaComputeContext::NextId()
	{
		if (sets->empty())
		{
			_currentId = 0;
		}
		else
		{
			_currentId += 1;
		}

		return _currentId;
	}

	ESGSet* CudaComputeContext::GetSet(int setId) const
	{

		vector<tuple<int, ESGSet*>>::iterator iterator = sets->begin();
		for (; iterator != sets->end(); ++iterator)
		{
			if (get<0>(*iterator) == setId)
				return get<1>(*iterator);
		}
		return nullptr;
	}

	operation_errors_t CudaComputeContext::ResetStatistics() const
	{
		if (memoryStatistics == nullptr)
			return esgSuccess;

		if (memoryStatistics->operationsList != nullptr)
		{
			memoryStatistics->operationsList.reset(new vector<unique_ptr<MemoryOperation>>);
		}
		return operation_errors_t::success();
	}

	operation_errors_t CudaComputeContext::SetCombinationModifiers(int setId, int matrixId, int weightsCount, float* weights)
	{
		cudaGetLastError(); //reset errors;
		if (weightsCount <= 0 || weights == nullptr)
			return operation_errors_t(esgInvalidParameters);


		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (weightsCount != (set->indicatorsCount) * (set->indicatorsCount))
			return operation_errors_t(esgParametersMismatch);

		if (set->GetMatrix(matrixId) == nullptr)
			return operation_errors_t(esgMatrixNotFound);

		float* d_matrix = nullptr;

		cudaError_t cudaStatusWeights = cudaMalloc((void **)&d_matrix, weightsCount * sizeof(float));

		this->AddMemoryOperation("cuda - matrix", weightsCount * sizeof(float));
		returnOnErrorMacro(cudaStatusWeights);

		cudaStatusWeights = cudaMemcpy(d_matrix, weights, weightsCount * sizeof(float), cudaMemcpyHostToDevice);
		if (cudaStatusWeights != cudaSuccess)
		{
			cudaFree(d_matrix);
			returnOnErrorMacro(cudaStatusWeights);
		}

		set->AddCombinationModifiers(matrixId, d_matrix);

		return operation_errors_t::success();
	}

	operation_errors_t CudaComputeContext::SetIssueWeights(int setId, int matrixId, int weightsCount, float* weights)
	{
		cudaGetLastError(); //reset errors;
		if (weightsCount <= 0 || weights == nullptr)
			return operation_errors_t(esgInvalidParameters);


		ESGSet* set = GetSet(setId);

		if (set == nullptr)
			return operation_errors_t(esgSetNotFound);

		if (weightsCount != (set->companiesCount) * (set->DISJUNCT_CLUSTERS_COUNT))
			return operation_errors_t(esgParametersMismatch);

		if (set->GetMatrix(matrixId) == nullptr)
			return operation_errors_t(esgMatrixNotFound);

		float* d_matrix = nullptr;

		cudaError_t cudaStatusWeights = cudaMalloc((void **)&d_matrix, weightsCount * sizeof(float));

		this->AddMemoryOperation("cuda - matrix", weightsCount * sizeof(float));
		returnOnErrorMacro(cudaStatusWeights);

		cudaStatusWeights = cudaMemcpy(d_matrix, weights, weightsCount * sizeof(float), cudaMemcpyHostToDevice);
		if (cudaStatusWeights != cudaSuccess)
		{
			cudaFree(d_matrix);
			returnOnErrorMacro(cudaStatusWeights);
		}

		set->AddIssuesWeights(matrixId, d_matrix);

		return operation_errors_t::success();
	}

	operation_errors_t CudaComputeContext::GovernanceRatings(int setId, int weightMatrixId, bool includeWeights, bool includeWeightedScores, GovernanceRatingsResult& result)
	{
		cudaProfilerStart();
		cudaError_t cudaStatus;

		cudaGetLastError(); //reset errors;

		this->ResetStatistics();
		ESGSet* set = GetSet(setId);

		if (set == nullptr)
		{
			return operation_errors_t(esgSetNotFound);
		}

		if (set->companiesCount != result.companiesCount)
			return operation_errors_t(esgCompaniesCountMismatch, "");

		if (set->indicatorsCount != result.indicatorsCount)
			return operation_errors_t(esgIndicatorsCountMismatch);


		float* weightsD = set->GetMatrix(weightMatrixId);
		if (weightsD == nullptr)
			return operation_errors_t(esgIndicatorsScoresNotFound);

		float* scoresD = set->indicators_scores;
		if (scoresD == nullptr)
			return operation_errors_t(esgMatrixNotFound);


		float* key_issues_weights = set->GetIssuesWeights(weightMatrixId);
		if (key_issues_weights == nullptr)
			return operation_errors_t(governanceWeightIssuesNotFound);

		float* modifiers = set->GetCombinationModifier(weightMatrixId);
		if (modifiers == nullptr)
			return operation_errors_t(governanceCombinationModifiersNotFound);

		int* companyAttributesD = set->company_flags;

		if (companyAttributesD == nullptr)
			return operation_errors_t(esgAttributesNotFound);


		int companiesNo = set->companiesCount;
		int indicatorsNo = set->indicatorsCount;
		int clustersCount = set->CLUSTERS_COUNT;
		
		int disjunctClusterCount = set->DISJUNCT_CLUSTERS_COUNT;
		int companyGroupTypeCount = result.groupingCategoriesCount;
		int* flagTypeCount = set->company_flag_types;
		int ratingsCount = companiesNo * (clustersCount);

		try
		{
			_mgpuContext->Start();

			cudaStatus = cudaMemcpyToSymbol(clusterStarts, set->h_configDisjunctClustersHeads, sizeof(byte) * set->h_configDisjunctClustersHeadsCount);
			returnOnErrorMacro(cudaStatus);
			cudaStatus = cudaMemcpyToSymbol(clusterGroups, set->h_configDisjunctClustersMap, sizeof(byte) * set->h_configDisjunctClustersMapCount);
			returnOnErrorMacro(cudaStatus);

			float* d_ratings;
			

			cudaStatus = cudaMalloc((void**)&d_ratings, sizeof(float) * ratingsCount);
			returnOnErrorMacro(cudaStatus);



			float* d_issues_Weighted_Scores;
			float* key_issues_weighted_scores_ptr = nullptr;

			if ( includeWeightedScores)
			{
				cudaMallocHost((void**)&d_issues_Weighted_Scores, sizeof(float) * ratingsCount);//throws CudaException
				key_issues_weighted_scores_ptr = d_issues_Weighted_Scores;
			}
			
			float* d_issues_Weights;
			float* d_issues_Weights_ptr = nullptr;

			if ( includeWeightedScores)
			{
				cudaMallocHost((void**)&d_issues_Weights, sizeof(float) * ratingsCount);//throws CudaException
				d_issues_Weights_ptr = d_issues_Weights;
			}

			esgKernels::ComputeRatings(clustersCount, weightsD, scoresD, modifiers, key_issues_weights, d_ratings, key_issues_weighted_scores_ptr, d_issues_Weights_ptr, companiesNo, indicatorsNo, disjunctClusterCount);

			cudaStatus = cudaPeekAtLastError(); //launch errors;
			returnOnErrorMacro(cudaStatus);


			cudaStatus = cudaMemcpy(result.ratings, d_ratings, ratingsCount * sizeof(float), cudaMemcpyDeviceToHost);
			returnOnErrorMacro(cudaStatus);


			esgKernels::CompanyAttributes* repeatedCompanyAttributes;
			int* d_benchmarking_maps;
			int rankingCount = companiesNo*clustersCount*companyGroupTypeCount;

			esgKernels::RepeatCompanyGroupFlags(companiesNo, companyAttributesD, d_ratings, clustersCount, companyGroupTypeCount, flagTypeCount, repeatedCompanyAttributes, d_benchmarking_maps);

			mgpu::MergesortPairs(repeatedCompanyAttributes, d_benchmarking_maps, rankingCount, CompanyAttributesComparer(), *_mgpuContext);

			cudaStatus = cudaMemcpy(result.indices, d_benchmarking_maps, rankingCount * sizeof(int), cudaMemcpyDeviceToHost);

			esgKernels::ComputeRanking(rankingCount, repeatedCompanyAttributes, result.rankings, result.percentile, result.group_counts);
			

			

			if ( includeWeights && (result.weights != nullptr))
			{
				cudaStatus = cudaMemcpy(result.weights, d_issues_Weights_ptr, ratingsCount * sizeof(float), cudaMemcpyDeviceToHost);
				if (cudaStatus != cudaSuccess)
				{
					returnOnErrorMacro(cudaStatus);
				}
			}

			if ( includeWeightedScores && (result.issue_Weighted_Scores != nullptr))
			{
				cudaStatus = cudaMemcpy(result.issue_Weighted_Scores, d_issues_Weighted_Scores, ratingsCount * sizeof(float), cudaMemcpyDeviceToHost);
				if (cudaStatus != cudaSuccess)
				{
					returnOnErrorMacro(cudaStatus);
				}
			}


			cudaMemset(clusterStarts, 0, sizeof(int) * MAX_DISJUNCT_CLUSTERS_HEADS);
			cudaMemset(clusterGroups, 0, sizeof(int) * MAX_DISJUNCT_CLUSTERS_GROUPS_MAP_COUNT);


			cudaFree(d_ratings);
			cudaFree(d_benchmarking_maps);
		}
		// ReSharper disable once CppEntityNeverUsed
		catch (mgpu::MgpuArchVersionException& exception)
		{
			return operation_errors_t(esgCudaInvalidArch, "MgpuArchVersion: this CUDA executable was not compiled with support for current device arch");
		}
		catch (mgpu::CudaException& exception)
		{
			cudaError_t status = exception.error;
			returnOnErrorMacro(status);
		}
		catch (...)
		{
			return operation_errors_t(esgError);
		}

		cudaProfilerStop();
		return operation_errors_t::success();
	}
}
